// Program: Odleglosc edycyjna w OpenCLu
// Autor: Marcin Milewski
#include <oclUtils.h>
#include <iostream>
#include <fstream>
#include <ctime>

// Name of the file with the source code for the computation kernel
// *********************************************************************
const char* cSourceFile = "dictionary-kernel.cl";

// OpenCL Vars
cl_context cxGPUContext;         // OpenCL context
cl_command_queue cqCommandQueue; // OpenCL command que
cl_platform_id cpPlatform;       // OpenCL platform
cl_device_id cdDevice;           // OpenCL device
cl_program cpProgram;            // OpenCL program
cl_kernel ckKernel;              // OpenCL kernel
size_t szGlobalWorkSize;         // 1D var for Total # of work items
size_t szLocalWorkSize;          // 1D var for # of work items in the work group	
size_t szParmDataBytes;          // Byte size of context information
size_t szKernelLength;           // Byte size of kernel code
cl_int ciErr1, ciErr2,           // Error code var
       ciErr3, ciErr4, ciErr5;
char* cPathAndName = NULL;       // var for full paths to data, src, etc.
char* cSourceCL = NULL;          // Buffer to hold source for compilation 

// zmienne programu
const int MAX_L = 16;                      // Maksymalna dlugosc slowa (łącznie z 0 na końcu napisu)
const int slow_w_pliku = 795875;           // liczba slow w pliku
cl_mem cmDevDictionary;                    // slownik na karcie
cl_mem cmResult;                           // wynik na karcie
cl_mem cmWord;                             // slowo skopiowane na karte
int* result;                               // miejsce na wynik na hoscie
char* charDict;                            // miesjce na splaszczony slownik na hoscie
int charDict_size = slow_w_pliku * MAX_L;  // rozmiar splaszczonego slownika
clock_t start_time, final_time;            // zegarki do pomiaru czasu

// Forward Declarations
// *********************************************************************
void Cleanup (int iExitCode);

// Main function 
// *********************************************************************
int main(int argc, char **argv)
{
    // start logs 
    shrSetLogFileName ("dictionary.log");

    // set and log Global and Local work size dimensions
    szLocalWorkSize = 256;

    szGlobalWorkSize = shrRoundUp((int)szLocalWorkSize, slow_w_pliku);     // rounded up to the nearest multiple of the LocalWorkSize
    shrLog("Global Work Size \t\t= %u\nLocal Work Size \t\t= %u\n# of Work Groups \t\t= %u\n\n", 
           szGlobalWorkSize, szLocalWorkSize, (szGlobalWorkSize % szLocalWorkSize + szGlobalWorkSize/szLocalWorkSize)); 

    // Allocate and initialize host arrays 
    shrLog( "Allocate and Init Host Mem...\n"); 
    result = (int *)calloc(slow_w_pliku * 4, 1);

    // Get an OpenCL platform
    ciErr1 = clGetPlatformIDs(1, &cpPlatform, NULL);
    shrLog("clGetPlatformID...\n"); 
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clGetPlatformID, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Get the devices
    ciErr1 = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevice, NULL);
    shrLog("clGetDeviceIDs...\n"); 
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clGetDeviceIDs, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Create the context
    cxGPUContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErr1);
    shrLog("clCreateContext...\n"); 
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clCreateContext, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Create a command-queue
    cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevice, 0, &ciErr1);
    shrLog("clCreateCommandQueue...\n"); 
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clCreateCommandQueue, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Allocate the OpenCL buffer memory objects for source and result on the device GMEM
    cmDevDictionary = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY,  sizeof(char) * charDict_size, NULL, &ciErr1);
    cmWord          = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY,  sizeof(char) * 32,            NULL, &ciErr2);
    cmResult        = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, 4 * slow_w_pliku,             NULL, &ciErr3);
    shrLog("clCreateBuffer...\n"); 
    if ((ciErr1|ciErr2|ciErr3) != CL_SUCCESS) {
        shrLog("ERROR in clCreateBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Read the OpenCL kernel in from source file
    shrLog("oclLoadProgSource (%s)...\n", cSourceFile); 
    cPathAndName = shrFindFilePath(cSourceFile, argv[0]);
    cSourceCL = oclLoadProgSource(cPathAndName, "", &szKernelLength);

    // Create the program
    cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &ciErr1);
    shrLog("clCreateProgramWithSource...\n"); 
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clCreateProgramWithSource, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Build the program with 'mad' Optimization option
    #ifdef MAC
        const char* flags = "-cl-fast-relaxed-math -DMAC";
    #else
        const char* flags = "-cl-fast-relaxed-math";
    #endif
    ciErr1 = clBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);
    shrLog("clBuildProgram...\n"); 
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clBuildProgram, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Create the kernel
    ckKernel = clCreateKernel(cpProgram, "TheKernel", &ciErr1);
    shrLog("clCreateKernel (TheKernel)...\n"); 
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clCreateKernel, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }


    shrLog("Wczytywanie slownika z pliku\n");
    typedef std::vector<std::string> Slownik;
    std::ifstream infile(argv[1]);
    Slownik vSlowa(slow_w_pliku);
    for (int i=0; i < slow_w_pliku; ++i) {
        infile >> vSlowa.at(i);
    }
    // skopiuj slownik na karte
    shrLog("Kopiowanie slowanika na karte\n");
    shrLog("  charDict_size: %d\n", charDict_size);
    charDict = (char*)calloc(charDict_size, 1);
    shrLog("  Slow w pliku: %d\n", slow_w_pliku);
    shrLog("  Splaszczanie slownika\n");
    for (int i=0; i < slow_w_pliku; ++i) {
        int offset_in_dict = MAX_L * i;
        for (int c=0; c < vSlowa.at(i).length(); ++c) {
            charDict[offset_in_dict + c] = vSlowa.at(i).at(c);
        }
    }
    ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, cmDevDictionary, CL_FALSE, 0, charDict_size, charDict,    0, NULL, NULL);
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clEnqueueWriteBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    start_time = clock();

    // wysylanie slowa na karte
    shrLog("Kopiowanie na karte slowa do sprawdzenia\n");
    std::string stringword = argv[2];
    const char* cstringword = stringword.c_str();
    int wordlen = stringword.length();
    ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, cmWord,          CL_FALSE, 0,            32, cstringword, 0, NULL, NULL);
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clEnqueueWriteBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Set the Argument values
    shrLog("Przygotowanie kernela -- zdefiniowane argumentów\n");
    ciErr1  = clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void*)&cmDevDictionary);
    ciErr1 |= clSetKernelArg(ckKernel, 1, sizeof(cl_mem), (void*)&cmResult);
    ciErr1 |= clSetKernelArg(ckKernel, 2, sizeof(cl_mem), (void*)&cmWord);
    ciErr1 |= clSetKernelArg(ckKernel, 3, sizeof(cl_int), (void*)&wordlen);
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clSetKernelArg, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Launch (run) kernel
    shrLog("Uruchamianie kernela\n");
    ciErr1 = clEnqueueNDRangeKernel(cqCommandQueue, ckKernel, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL);
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clEnqueueNDRangeKernel, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    // Synchronous/blocking read of results, and check accumulated errors
    shrLog("Odczytywanie wyników z karty\n"); 
    ciErr1 = clEnqueueReadBuffer(cqCommandQueue, cmResult, CL_TRUE, 0, slow_w_pliku * sizeof(int), result, 0, NULL, NULL);
    if (ciErr1 != CL_SUCCESS) {
        shrLog("ERROR in clEnqueueReadBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
        Cleanup(EXIT_FAILURE);
    }

    final_time = clock();
    std::cout << "\nCzas wykonania (kernel + skopiowanie wyniku): " << 1000*(final_time - start_time)/(double)CLOCKS_PER_SEC << " ms\n" << std::endl;

    shrLog("Szukanie najlepszego dopasowania\n");
    int min_dist = 88888;
    int where_is_min = -1;
    for (int i=0; i<slow_w_pliku; ++i) {
        // if (i<100) std::cout << result[i] << " ";
        if (result[i] < min_dist) {
            min_dist = result[i];
            where_is_min = i;
        }
    }
    if (where_is_min > -1) {
        std::cout << "Closest match (at " << where_is_min << "): " << vSlowa.at(where_is_min) << "  distance: " << min_dist << std::endl;
    }

    // Cleanup and leave
    Cleanup (EXIT_SUCCESS);
}

void Cleanup (int iExitCode)
{
    // Cleanup allocated objects
    shrLog("Starting Cleanup...\n\n");
    if(cPathAndName)free(cPathAndName);
    if(cSourceCL)free(cSourceCL);
    if(ckKernel)clReleaseKernel(ckKernel);  
    if(cpProgram)clReleaseProgram(cpProgram);
    if(cqCommandQueue)clReleaseCommandQueue(cqCommandQueue);
    if(cxGPUContext)clReleaseContext(cxGPUContext);

    if(cmResult) clReleaseMemObject(cmResult);
    if(cmDevDictionary) clReleaseMemObject(cmDevDictionary);
    if(cmWord) clReleaseMemObject(cmWord);

    if(result) free(result);

    shrLog("Done, Exiting...");
    exit (iExitCode);
}
